/*
Copyright (c) [2019年5月1日] [吴超]
[MBT_studio] is licensed under Mulan PSL v2.
You can use this software according to the terms and conditions of the Mulan PSL v2.
You may obtain a copy of Mulan PSL v2 at:
		 http://license.coscl.org.cn/MulanPSL2
THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND,
EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT,
MERCHANTABILITY OR FIT FOR A PARTICULAR PURPOSE.
See the Mulan PSL v2 for more details.
*/
#include "build_AdjacentUnit.cuh"

#include "core/CUDA_utils.cuh"
#include <线性代数/矢量计算.h>



//填充体素遮罩
__global__ void f_Fill_VolexMask(uint8* 体素) {
	//uint64 id = blockIdx.x * blockDim.x + threadIdx.x;
	//uint64 id = threadIdx.x + (blockIdx.x * gridDim.x) + (blockDim.x * gridDim.x * blockIdx.y);
	uint64 id = (blockDim.x * gridDim.x * blockIdx.y) + (gridDim.x * blockIdx.x) + threadIdx.x;
	//uint32 id = blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
	体素[id] = 0;
	//printf("f_Fill_VolexMask = %i \n", id);
}



/*

__device__ inline uint32 f_CoordToHash(const float3& vert) {
	return uint32(vert.x + 511) | (uint32(vert.y + 511) << 10) | (uint32(vert.z + 511) << 20);
}*/

//单元分类
__global__ void f_UnitClassification(vec3* 顶点, uint32* keys, uint32* value, S_DomProp dom/* const uvec3 dim, const vec3 元点, float32 单元大小*/) {
	//uint32 id = threadIdx.x + (blockDim.x * (gridDim.x * blockIdx.y + blockIdx.x));
	//uint32 id = (threadIdx.z * blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;
	uint64 id = blockIdx.x * blockDim.x + threadIdx.x;
	//uint64 id = (blockIdx.y * gridDim.x) + blockIdx.x;
	//printf("f_UnitClassification = %i \n", id); 顶点[id].x; 顶点[id].y; 顶点[id].z;
	//return;
	//int3 gridPos;
	/*gridPos.x = floor((顶点[id].x - dom.m_域世界偏移.x) / dom.m_单元大小);
	gridPos.y = floor((顶点[id].y - dom.m_域世界偏移.y) / dom.m_单元大小);
	gridPos.z = floor((顶点[id].z - dom.m_域世界偏移.z) / dom.m_单元大小);*/
	ivec3 gridPos;
	gridPos.x = int32((顶点[id].x - dom.m_域世界偏移.x) / dom.m_单元大小);
	gridPos.y = int32((顶点[id].y - dom.m_域世界偏移.y) / dom.m_单元大小);
	gridPos.z = int32((顶点[id].z - dom.m_域世界偏移.z) / dom.m_单元大小);
	//printf("f_UnitClassification = %i %lli \n", id, (gridPos.z * dom.m_分辨率.x * dom.m_分辨率.y) + (gridPos.y * dom.m_分辨率.x) + gridPos.x);
	//printf("f_UnitClassification =  %lli %i %i %i \n", id, gridPos.x, gridPos.y, gridPos.z);
	/*vec3 vert = 顶点[id];
	vert.x /= 单元大小;
	vert.y /= 单元大小;
	vert.z /= 单元大小;*/
	//return;
	keys[id] = (gridPos.z * dom.m_分辨率.x * dom.m_分辨率.y) + (gridPos.y * dom.m_分辨率.x) + gridPos.x;
	value[id] = id;
	//keys[id].x = blockIdx.x;
	//keys[id].y = blockIdx.y;
	//printf("f_UnitClassification = %lli = %i %i %i \n", id, gridPos.x, gridPos.y, gridPos.z);

}




//查找索引开始位置
__global__ void f_FindIndexStart(uint32* keys, uint32* 体素, uint8* 索引数量, const uvec3 dim) {
	//uint32 id = (blockIdx.y * gridDim.x) + blockIdx.x;
	uint64 id = blockIdx.x * blockDim.x + threadIdx.x;

	if (id) {
		if (keys[id] != keys[++id]) {
			//printf("id = %i", id);
			//printf(" . . . . key %i", keys[id]);
			uint64 k = keys[id];

			//uvec3 coord = f_线性索引转坐标(k, dim);
			体素[k] = id;
			//*sliceHead = id;
			索引数量[k] = 0;
			for (uint8 i = 0; i < 128; ++i) {
				if (k == keys[id + i]) {
					++索引数量[k];
				}
			}
		}
	}
	else {
		uint64 k = keys[id];

		体素[k] = id;
		//*sliceHead = id;
		索引数量[k] = 0;
		for (uint8 i = 0; i < 64; ++i) {
			if (k == keys[id + i]) {
				++索引数量[k];
			}
		}
	}
	//uint32* sliceHead = (uint32*)((char*)体素.ptr + 体素.pitch * keys[id]);
	//uint32* rowHead = (uint32*)((char*)sliceHead + y * devPitchedPtr.pitch);
	
	//体素[id] = keys[id];
	//((uint32*)(体素.ptr))[id] = keys[id];
	//printf(" key =  %i", ((uint32*)(体素.ptr))[id]);
	//printf("\t");
}




void ouf_分配单元区间GPU内存(S_邻接数据& data) {
	//const uvec4& 分辨率 = data.m_域属性.m_单元属性.m_分辨率;

	if (data.m_邻接索引_Key != nullptr) cudaFree(data.m_邻接索引_Key);
	if (data.m_邻接索引_Value != nullptr) cudaFree(data.m_邻接索引_Value);
	if (data.m_体素索引 != nullptr) cudaFree(data.m_体素索引);
	if (data.m_单元内邻接数量 != nullptr) cudaFree(data.m_单元内邻接数量);
	

	COUDA_CALL(cudaMalloc((void**)&(data.m_邻接索引_Key), sizeof(uint32) * data.m_Num * 2));
	COUDA_CALL(cudaMalloc((void**)&(data.m_邻接索引_Value), sizeof(uint32) * data.m_Num * 2));
	COUDA_CALL(cudaMalloc((void**)&(data.m_体素索引), sizeof(uint32) * data.m_体素Num));

	COUDA_CALL(cudaMalloc((void**)&(data.m_单元内邻接数量), sizeof(uint8) * data.m_体素Num));
}




void cuf_分类邻接单元(S_邻接数据& data, vec3* loc, uint32 线程块大小) {
	const uvec3& 分辨率 = data.m_域属性.m_单元属性.m_分辨率;
	//const uvec3 Dim = ivec_4转3(分辨率);

	dim3 gridDim = { 1,1,1 };
	dim3 blockDim = { 1,1,1 };


	f_Fill_VolexMask << < {分辨率.y, 分辨率.z, 1}, { 分辨率.x, 1, 1} >> > (data.m_单元内邻接数量);
	cudaDeviceSynchronize();

	if (data.m_Num > 线程块大小) {
		blockDim.x = 线程块大小;
		gridDim.x = data.m_Num / 线程块大小;
	}
	else {
		blockDim.x = data.m_Num;
	}
	//f_Fill_IndexVolex << < gridDim, { 1 } >> > (data.m_体素索引, data.m_Dim);
	f_UnitClassification << < gridDim, blockDim >> > (loc, data.m_邻接索引_Key, data.m_邻接索引_Value, data.m_域属性.m_单元属性);
	cudaDeviceSynchronize();

	//排序键值
	thrust::device_ptr<uint32> keys(data.m_邻接索引_Key);
	thrust::device_ptr<uint32> value(data.m_邻接索引_Value);
	thrust::sort_by_key(keys, keys + data.m_Num, value);
	cudaDeviceSynchronize();
	
	//统计单元索引数量 在体素索引中标记索引
	f_FindIndexStart <<< gridDim, blockDim >>> (data.m_邻接索引_Key, data.m_体素索引, data.m_单元内邻接数量, data.m_域属性.m_单元属性.m_分辨率);
	cudaDeviceSynchronize();
	return;
 }


void buf_数据拷贝回内存(uint32* index, uint8* iNum, uint32* v, S_邻接数据& data) {
	const uvec3& 分辨率 = data.m_域属性.m_单元属性.m_分辨率;
	uint32 num = 分辨率.x * 分辨率.y * 分辨率.z;

	cudaMemcpy(v, data.m_体素索引, num * sizeof(uint32), cudaMemcpyKind::cudaMemcpyDeviceToHost);
	cudaMemcpy(iNum, data.m_单元内邻接数量, num * sizeof(uint8), cudaMemcpyKind::cudaMemcpyDeviceToHost);

	cudaMemcpy(index, data.m_邻接索引_Value, data.m_Num * sizeof(uint32), cudaMemcpyKind::cudaMemcpyDeviceToHost);
	
}









/*void cuf_排序邻接单元(S_邻接数据& data, vec3* loc) {
	cudaPitchedPtr GPU3D数组;
	f_填充索引体素 <<< 1, 1 >>> (data.m_体素索引, data.m_Dim);
}*/











